#include <Particula.h>

__global__ void CalculaIDGridCell(float3* vetorPosicoes, unsigned int* values, unsigned int* keys, int numeroParticulas, int numColunas, float deltaX, float deltaY, unsigned int* count, float XminGrid, float YminGrid) {

    int thread = threadIdx.x;
    int block = blockIdx.x;
    int dim = blockDim.x;
    int indiceValue = dim*block + thread;
    int idCell;

    if (indiceValue < numeroParticulas) {
        int posicaoParticula = values[indiceValue];
        int j = (vetorPosicoes[posicaoParticula].x) / deltaX;
        int i = (vetorPosicoes[posicaoParticula].y) / deltaY;
        idCell = (i * numColunas) + j;
        keys[indiceValue] = idCell;
        atomicInc(&count[idCell], numeroParticulas);
    }
}

extern "C" void launch_CalculaIDGridCell(unsigned int blocksPerGrid, unsigned int threadsPerBlock,
        float3* posicoes, unsigned int* values, unsigned int* keys, int numeroParticulas,
        int numColunas, float deltaX, float deltaY, unsigned int* count, float XminGrid, float YminGrid) {
    // execute the kernel
    CalculaIDGridCell <<<blocksPerGrid, threadsPerBlock>>>(posicoes, values, keys, numeroParticulas, numColunas, deltaX, deltaY, count, XminGrid, YminGrid);
}

__global__ void ZeraCountEOffSet(int numeroCelulas, unsigned int* count, unsigned int* offSet) {

    int thread = threadIdx.x;
    int block = blockIdx.x;
    int dim = blockDim.x;
    int indice = dim*block + thread;

    if (indice < numeroCelulas) {
        count[indice] = 0;
        offSet[indice] = 0;
    }
}

extern "C" void launch_ZeraCountEOffSet(unsigned int blocksPerGrid, unsigned int threadsPerBlock, int numeroCelulas,
        unsigned int* count, unsigned int* offSet) {
    // execute the kernel
    ZeraCountEOffSet <<<blocksPerGrid, threadsPerBlock>>>(numeroCelulas, count, offSet);
}
